feat(KV-FP8): the CUDA fp8 KV store and read, reached by removing the W1 guard that refused CUDA before the provider table (#1593) - #1606
Merged
Conversation
… W1 guard that refused CUDA before the provider table (#1593) `KV-FP8` W1 landed the CPU half in 2026-07 and left the CUDA arm as a named later brick. It is now the critical path of benchmark campaign #1574, whose subject `r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121` declares `kv_cache_quant_algo: "FP8"` in `hf_quant_config.json`, so no cell of that three-way can be served correctly without it. W1 is the ORACLE for this wave. Every gate here compares CUDA to the landed CPU kernels; nothing re-ports the numerics. THE STORE is a 1:1 port of the fp8 branch of vLLM's `reshape_and_cache_flash_kernel` (`csrc/libtorch_stable/cache_kernels.cu:314-401`) plus `CopyWithScaleOp` (`:241-252`) at pin `555967922`, restricted to upstream's `is_contiguous_heads && kv_scale_stride == 0` arm (`:352-366`) — the only arm the op's wrapper admits, because the vt cache is the NHD unbind slice and `ReshapeAndCacheFp8` takes two scalar scales. The converter is upstream's own `__nv_cvt_float_to_fp8(hp / scale, __NV_SATFINITE, __NV_E4M3)` (`quant_utils.cuh:497-503`), a true DIVIDE rather than the activation path's hoisted reciprocal multiply, and its byte-for-byte equality to the CPU software codec `vt::F32ToF8E4M3` is already MEASURED at zero tolerance on sm_110 and sm_121a (`.agents/specs/vt-fp8-quant-arch-gate.md` G2). THE READ adds `LoadKv(ptr, i, scale)` beside `Load`. It is INERT on the f32 and bf16 arms — they forward to `Load` unchanged, so every existing caller reads the same bytes in the same order — and on `uint8_t` it is upstream's `scaled_vec_conversion<float, uint8_t>` (`quant_utils.cuh:302-308`), written as the SAME ARITHMETIC as `vt::F8E4M3ToF32` so that CUDA==CPU on the read is a property of the source rather than of a measurement. Only the two correctness-grade kernels serve fp8, which is what the existing ladder already implies: the WMMA prefill kernels stage bf16 fragments, the vendored FA-2 launchers take bf16 pointers, and the vectorized decode-opt/GQA kernels read through `LoadRowN`/`LoadRow8`, 128-bit `uint4` loads specialized for bf16 and f32 only. Upstream draws the same line from the other side (`flash_attn.py:181-187,796-805`). WHAT ACTUALLY MADE THE ARM UNREACHABLE was neither kernel. Both W1 wrappers carried `VT_CHECK(q.device.type == DeviceType::kCPU, ... "a named later brick")` evaluated BEFORE provider lookup, so no CUDA kernel could ever have been reached however well it was registered. That is the RED this change was written against. The STORE now resolves through the provider table like every other op, because `kReshapeAndCacheFp8` is its own `OpId` that only CPU and CUDA register and an unimplemented backend refuses BY NAME inside `GetOp`. The READ cannot: it rides ADDITIVE fields on `PagedAttentionArgs` of an op `kMETAL` and `kROCM` already register for the FLOAT path, and nothing in the provider table separates the two arms, so an fp8 cache would reach a float kernel and return silent garbage. It keeps an explicit CPU-or-CUDA list whose message names the missing part. UNREACHED, DELIBERATELY, AND NOT NEW. Nothing calls the fp8 KV path from a production entry point on either backend: `vt::ReshapeAndCacheFp8` and `PagedAttentionArgs::kv_cache_dtype` have no caller outside their tests. W1 landed in that state and this does not change it. `KV-FP8` W3 owns the wiring — half-sized KV blocks in the runner, `--kv-cache-dtype` threaded from the CLI, and the checkpoint `k_scale`/`v_scale` path — and #1593 tracks it. Listed under `## Owed` in `.agents/specs/fp8-kv-cache.md`. THE DEVICE GATES ARE UNEXECUTED AND THE CUDA TRANSLATION UNITS ARE UNCOMPILED. The implementing session had no CUDA toolkit (`nvcc` is absent on `mudler-ubuntu-box`) and no device (the fleet was leased for #1574), so G2 (provider registration), G3 (store byte parity, f32 and bf16), G4 (paged-read parity, decode and prefill) and G5 (the e5m2 refusal) all skip with a MESSAGE naming what did not run. G1 and G1b — provider routing and the Metal/ROCm refusal — are the only cases that ran, and they run on the CPU leg. The first CUDA build or `rc` lease that touches this row must run `ctest -R test_cuda_fp8_kv_cache` before W2 counts as measured; the spec's `## Owed` says so. EVIDENCE. RED first: `test_cuda_fp8_kv_cache` 6 assertions / 6 failed, on `reshape_and_cache_fp8: only the CPU fp8-KV store is implemented in W1` at `ops.cpp:3478` and `paged_attention: only the CPU fp8-KV read is implemented in W1` at `ops.cpp:3811`. Green after: 6 cases / 10 assertions. Three negative mutations, each rebuilt, run and restored against a pre-taken sha256 — reinstate the store guard (3 failed), reinstate the read guard (7 failed), delete the Metal/ROCm refusal (4 failed). No sibling regressions: `test_ops_fp8_kv_cache` 8/511, `test_ops_reshape_cache` 12/192, `test_ops_paged_attn` 14/1646, `test_ops_paged_attn_dtype` 3/172. Issue: #1593 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…ld not fail for any CUDA defect (#1593, #1636) Repairs the six findings of the fresh review of #1593 W2. Four are records, two are gates, and the two merge-blocking ones land in the commit message, which `squash_merge_commit_message = PR_BODY` makes permanent. THE ANCHOR. Four new sites and the pull request body cited `scaled_vec_conversion<float, uint8_t>` at `quant_utils.cuh:302-308`. Verified against the pinned tree at `5559679229bc961848b121ccdeaa8fa5d79bec98`: lines 301-305 are the GENERIC primary template, whose body is `return x;`, and 307-314 are the `// fp8 -> half` `<uint16_t, uint8_t>` specialization. The `<float, uint8_t>` one is at `:419-429`. Same class, same change: `Fp8KVCacheDataType` was cited at `dtype_fp8.cuh:9-13`, the `#include <cuda_fp8.h>` guard; the enum is at `:15-19`. An anchor is how the next reader checks a port against the oracle, and one landing on an identity template invites the conclusion that the port is unfaithful. Three W1 copies of the same wrong anchor (`include/vt/fp8_kv.h:92`, `include/vt/ops.h:1129`, `src/vt/cpu/cpu_paged_attn.cpp:164`) are outside this change's authority and are filed as #1636, owned by `KV-FP8` and listed under the spec's `## Owed`. THE TRANSLATION UNITS COMPILE. The body said they do not. CI job `cuda-fat-build` built both changed files for `80;86;87;89;90a;100a;103a;110; 120a;121a` under `-Werror=all-warnings` and PASSED on `4d71e776e`, run 32495320287, job 96812232428. What stays true is that nothing has been EXECUTED on a device, because that job configures `-DVLLM_CPP_BUILD_TESTS=OFF`. The spec now separates the two states. G5 COULD NOT FAIL FOR ANY CUDA DEFECT, and the fix needed three mutations to state correctly. The e5m2 refusal exists in the op wrapper, the CPU kernel and the CUDA kernel. Deleting the CUDA one on a CPU build gives `ninja: no work to do` and leaves the file 7/10 SUCCESS. Deleting the wrapper's leaves W1's `test_ops_fp8_kv_cache` GREEN at 8/511, because execution falls through to the CPU kernel's check and W1's case asserts `CHECK_THROWS_AS` on `std::runtime_error` rather than a message; only deleting both turns it red (7/8, 510/511). A layered refusal needs an assertion that names its layer, so G5 now resolves the registered provider with `GetOp(OpId::kReshapeAndCacheFp8, DeviceType::kCUDA)`, calls it directly, and requires both `cuda reshape_and_cache_fp8` and `fp8_e5m2` in the message. TWO UNGATED INSTANTIATIONS. `LaunchPagedFp8Out<__nv_bfloat16, __nv_bfloat16>` is what a served bf16 model takes and what #1574's subject runs, and G4 exercised only `<float, float>`; the store's `DType::kF16 -> __half` arm was untouched while the wrapper's `IsFloat()` admits it. G4b and the f16 leg of G3 close both. Neither has a red-first mutation: they are device cases and this session has no `nvcc` and no device, which the first mutation above measures rather than assumes. They are listed under `## Owed`. TWO PROSE OVERSTATEMENTS. The store is elementwise-identical, not 1:1: upstream vectorizes the contiguous-heads arm through `vectorize_with_alignment<VEC_SIZE>` (`cache_kernels.cu:360-363`) and this is a scalar strided loop over the same elements in the same order, which is a bandwidth difference W4 owns. And the read is the CPU codec for every one of the 254 finite e4m3 codes, not "line for line": on `0x7F` and `0xFF` the CPU returns `quiet_NaN()` (`0x7FC00000`) and the device returns `CUDART_NAN_F` (`0x7FFFFFFF`), which no gate here can see because a NaN compares unequal to itself. Focused gate after: `test_cuda_fp8_kv_cache` 7 cases / 10 assertions SUCCESS, Release CPU build with `-Wall -Wextra -Werror`. Siblings unmoved: `test_ops_fp8_kv_cache` 8/511, `test_ops_reshape_cache` 12/192, `test_ops_paged_attn` 14/1646, `test_ops_paged_attn_dtype` 3/172. Every mutation restored against a pre-taken sha256. Issue: #1593 Anchor debt: #1636 FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
…at CI falsified Both matrix rows said the W2 CUDA translation units are uncompiled. That was true of the implementing session and false of this branch: CI `cuda-fat-build` built them for ten architectures under `-Werror=all-warnings` on `4d71e776e` (run 32495320287). What is still true, and now stated instead, is that NOTHING HAS EXECUTED them, because that job configures `-DVLLM_CPP_BUILD_TESTS=OFF`. Those two clauses were outside the repairing implementer's authority, so they were tracked in #1636 and left contradicting the spec they link to. A record edit rides in the pull request whose change made the record stale, so they ride here. The three W1-side copies of the wrong `quant_utils.cuh` anchor stay in #1636, because W1 landed separately and this branch did not stale them. FOLLOWING_AGENTS_PROTOCOL Following-Agents-Protocol: true AI-Assisted: true Assisted-by: AGENT:claude-opus-5 [Claude Code]
localai-bot
marked this pull request as ready for review
August 21, 2026 20:39
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
feat(KV-FP8): the CUDA fp8 KV store and read, reached by removing the W1 guard that refused CUDA before the provider table (#1593)
KV-FP8W1 landed the CPU half in 2026-07 and left the CUDA arm as a namedlater brick. It is now the critical path of benchmark campaign #1574, whose
subject
r0b0tlab/Qwen3.8-27B-NVFP4-MTP-sm121declareskv_cache_quant_algo: "FP8"inhf_quant_config.json, so no cell of thatthree-way can be served correctly without it.
W1 is the ORACLE for this wave. Every gate here compares CUDA to the landed CPU
kernels. Nothing re-ports the numerics.
The store
An ELEMENTWISE-IDENTICAL port of the fp8 branch of vLLM's
reshape_and_cache_flash_kernel(csrc/libtorch_stable/cache_kernels.cu:314-401)plus
CopyWithScaleOp(:241-252) at pin555967922, restricted to upstream'sis_contiguous_heads && kv_scale_stride == 0arm (:352-366). That is the onlyarm the op's wrapper admits, because the vt cache is the NHD unbind slice and
ReshapeAndCacheFp8takes two scalar scales.It is not instruction-identical, and calling it a 1:1 port would overstate it.
Upstream's contiguous-heads arm moves the row through
vectorize_with_alignment<VEC_SIZE>(:360-363,VEC_SIZE8 for a 2-bytesource and 4 for f32); this one is a scalar strided loop over the same elements
in the same order. Same bytes out, fewer bytes per instruction. W4 owns the
bandwidth.
The converter is upstream's own
__nv_cvt_float_to_fp8(hp / scale, __NV_SATFINITE, __NV_E4M3)(
quant_utils.cuh:497-503). It is a true DIVIDE, not the activation path'shoisted reciprocal multiply, and the two differ by up to one f32 ulp before the
round. Its byte-for-byte equality to the CPU software codec
vt::F32ToF8E4M3isalready MEASURED at zero tolerance on sm_110 and sm_121a
(
.agents/specs/vt-fp8-quant-arch-gate.mdG2).The read
LoadKv(ptr, i, scale)joinsLoad. It is INERT on the f32 and bf16 arms, whichforward to
Loadunchanged, so every existing caller reads the same bytes in thesame order. On
uint8_tit is upstream'sscaled_vec_conversion<float, uint8_t>(
quant_utils.cuh:419-429), written as the SAME ARITHMETIC asvt::F8E4M3ToF32so that CUDA equals CPU on the read as a property of the source rather than of a
measurement.
Same arithmetic holds for all 254 FINITE e4m3 codes and not for the two NaN ones.
On
0x7Fand0xFFthe CPU returnsstd::numeric_limits<float>::quiet_NaN()(
0x7FC00000) and the device returnsCUDART_NAN_F(0x7FFFFFFF): both quiet,both propagating, different payload. No gate here can see it, because a NaN
compares unequal to itself, so a byte or NMSE comparison fails on any payload
rather than on the wrong one. Reaching it also needs a non-finite input, since
__NV_SATFINITEclamps an out-of-range magnitude to0x7Eor0xFE.Only the two correctness-grade kernels serve fp8, which is what the existing
ladder already implies. The WMMA prefill kernels stage bf16 fragments, the
vendored FA-2 launchers take bf16 pointers, and the vectorized decode-opt and GQA
kernels read through
LoadRowNandLoadRow8, 128-bituint4loads specializedfor bf16 and f32 only. Upstream draws the same line from the other side, at
flash_attn.py:181-187,796-805. A tensor-core fp8 read is a performance brick;W2's gate is parity and W4 owns the memory measurement.
What actually made the arm unreachable
Neither kernel. Both W1 wrappers carried
VT_CHECK(q.device.type == DeviceType::kCPU, ... "a named later brick")evaluated BEFORE provider lookup, so no CUDA kernel could ever have been reached
however well it was registered. That is the RED this change was written against.
The STORE now resolves through the provider table like every other op, because
kReshapeAndCacheFp8is its ownOpIdthat only CPU and CUDA register, so anunimplemented backend refuses by name inside
GetOp. The READ cannot: it ridesADDITIVE fields on
PagedAttentionArgsof an opkMETALandkROCMalreadyregister for the FLOAT path, and nothing in the provider table separates the two
arms, so an fp8 cache would reach a float kernel and return silent garbage. It
keeps an explicit CPU-or-CUDA list whose message names the missing part.
Unreached, deliberately, and not new
Nothing calls the fp8 KV path from a production entry point on either backend.
vt::ReshapeAndCacheFp8andPagedAttentionArgs::kv_cache_dtypehave no calleroutside their tests. W1 landed in that state and this does not change it.
KV-FP8W3 owns the wiring: half-sized KV blocks in the runner,--kv-cache-dtypethreaded from the CLI, and the checkpointk_scaleandv_scalepath. #1593 tracks it, and.agents/specs/fp8-kv-cache.mdlists itunder
## Owed.The CUDA translation units COMPILE; the device gates are UNEXECUTED
Those are two different states and the first version of this body ran them
together. CI job
cuda-fat-build(.github/workflows/ci.yml:773) configures-DVLLM_CPP_CUDA=ONwith-DVLLM_CPP_CUDA_ARCHITECTURES='80;86;87;89;90a;100a;103a;110;120a;121a', buildstarget
vllm, and reaches every CUDA translation unit with-Werror=all-warnings. Both files this change touches are unconditional sourcesof that target. Line anchors as read at the tree CI built,
4d71e776efc18cb5e61a26e642ddad8de5339134, and not at this branch's merge base:cmake/CompilerWarnings.cmake:53,CMakeLists.txt:1694and:1697. That jobPASSED on that same SHA: run 32495320287, job 96812232428.
So the CUDA arm was uncompiled by the IMPLEMENTING session, which had no
nvccand no device, and has since been compiled for ten architectures.
Execution is still absent, because that same job configures
-DVLLM_CPP_BUILD_TESTS=OFFand therefore never links or runstests/vt/test_cuda_fp8_kv_cache.cpp. G2 (provider registration), G3 (store byteparity over the f32, bf16 and f16 sources), G4 and G4b (paged-read parity over
the decode and prefill shapes, for the f32 and the bf16 query/output
instantiation) and G5 (the CUDA kernel's own e5m2 refusal) all skip with a
MESSAGEnaming what did not run. G1 and G1b, provider routing and theMetal/ROCm refusal, are the only cases that ran, and they run on the CPU leg.
A CPU-only gate cannot see a CUDA defect, and the fresh review PROVED this one
blind rather than arguing it: deleting the entire production call site for the
CUDA fp8 read left the focused gate 100% green, and inverting the CUDA store's
scale direction produced
ninja: no work to do, because a CPU build'scompile_commands.jsoncontains ZERO.cutranslation units — 1046 entries inthe review's configure, 1030 in this session's, not one of them CUDA. That
is the honest consequence of the state above. The first
rclease that touchesthis row must run
ctest -R test_cuda_fp8_kv_cachebefore W2 counts as measured.The spec's
## Owedsays so, and the wave table'sDONEmeans landed and gatedrather than measured on hardware.
What the fresh review changed
Six findings, all repaired here.
Four new sites plus this body cited
scaled_vec_conversion<float, uint8_t>atquant_utils.cuh:302-308. At the pin,301-305is the GENERIC primary templatethat returns its input unchanged and
307-314is the fp8-to-HALF<uint16_t, uint8_t>specialization. The<float, uint8_t>one is at:419-429.Fp8KVCacheDataTypewas cited atdtype_fp8.cuh:9-13, the#include <cuda_fp8.h>guard; the enum is at:15-19. Three W1 copies of thesame wrong anchor are outside this change's authority and are filed as #1636.
G5 could not fail for any CUDA-side defect. It asserted a bare throw that
originates in the op wrapper, ABOVE the device checks and above
GetOp. Thee5m2 refusal exists in THREE places — that wrapper, the CPU kernel and the CUDA
kernel — and both halves of the finding were measured rather than argued.
Deleting the CUDA kernel's
VT_CHECKon a CPU build givesninja: no work to doand leaves the file 7 cases / 10 assertionsSUCCESS!. Deleting thewrapper's leaves W1's
test_ops_fp8_kv_cacheGREEN at 8/511, because executionfalls through to the CPU kernel's own check and W1's
refuses e5m2case(
tests/vt/test_ops_fp8_kv_cache.cpp:342) assertsCHECK_THROWS_ASonstd::runtime_errorrather than a message; only deleting the wrapper AND theCPU kernel check turns it red, at 7/8 cases and 510/511 assertions. So what W1
pins is "e5m2 is refused somewhere on the CPU path", and a layered refusal needs
an assertion that NAMES its layer. G5 now resolves the registered provider with
GetOp(OpId::kReshapeAndCacheFp8, DeviceType::kCUDA), calls it directly, andrequires the message to carry both
cuda reshape_and_cache_fp8andfp8_e5m2,which no other layer produces. Three mutations, each rebuilt, run and restored
against a pre-taken sha256.
LaunchPagedFp8Out<__nv_bfloat16, __nv_bfloat16>— the instantiation a servedbf16 model takes, and the one #1574's subject runs — was entirely ungated, and
the store's
DType::kF16arm was untouched while the wrapper'sIsFloat()admits it. G4b and the f16 leg of G3 close both.
Two prose overstatements are now true: the store is elementwise-identical rather
than 1:1, and the read is the CPU codec for every finite code rather than "line
for line".
Evidence
RED first, in the implementing session:
test_cuda_fp8_kv_cachereported 6assertions and 6 failed, on
reshape_and_cache_fp8: only the CPU fp8-KV store is implemented in W1atops.cpp:3478andpaged_attention: only the CPU fp8-KV read is implemented in W1atops.cpp:3811.Three negative mutations from that session, each rebuilt, run and restored
against a pre-taken sha256. Reinstating the store guard failed 3 assertions.
Reinstating the read guard failed 7. Deleting the Metal/ROCm refusal failed 4.
GREEN AFTER, on the repaired tree:
test_cuda_fp8_kv_cache7 cases / 10assertions /
SUCCESS!, exit 0, on a Release CPU build with-Wall -Wextra -Werror. It was 6 cases before; the extra case is G4b, and it skips here withthe other four device cases, each printing a
MESSAGEnaming what did not run.No sibling regressions, same tree:
test_ops_fp8_kv_cache8 cases / 511assertions,
test_ops_reshape_cache12 / 192,test_ops_paged_attn14 / 1646,test_ops_paged_attn_dtype3 / 172, allSUCCESS!.THE THREE CASES ADDED OR REWRITTEN HERE HAVE NO RED-FIRST MUTATION. G3's f16
leg, G4b and the rewritten G5 are device cases, this session has no
nvccand nodevice for the same reason the implementing session had none, and that a CPU
build can neither compile nor execute them is measured above rather than assumed.
They are listed
under
## Owedin.agents/specs/fp8-kv-cache.mdwith the rest of the devicehalf, and the
rclease that runsctest -R test_cuda_fp8_kv_cacheowes theirmutations too.
Issue: #1593
Anchor debt filed: #1636
FOLLOWING_AGENTS_PROTOCOL
Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [Claude Code]